#include "common.hpp"
#include "utility.hpp"

namespace cuda
{
////////////////////////////////////////////////////////////////////////////////////////////////////
__global__ 
void __add(
	float2* a,			
	float2  b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < nRows*nCols)
	{
		c[idx].x = a[idx].x + b.x;
		c[idx].y = a[idx].y + b.y;
	}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
void add(
	float2* a,			
	float2  b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	__add config(nRows*nCols, 1, 256, 1, 0)	\
		(a, b, c, nRows, nCols);
	cudaThreadSynchronize();
	cudaGetLastError();
}

////////////////////////////////////////////////////////////////////////////////////////////////////
__global__ 
void __add(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < nRows*nCols)
	{
		c[idx].x = a[idx].x + b[idx].x;
		c[idx].y = a[idx].y + b[idx].y;
	}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
void add(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	float2* tmp;
	cudaMalloc((void**)&tmp, sizeof(float2)*nCols*nRows);

	__add config(nRows*nCols, 1, 256, 1, 0)	\
		(a, b, tmp, nRows, nCols);
	cudaMemcpy(c, tmp, sizeof(float2)*nCols*nRows, cudaMemcpyDeviceToDevice);
	cudaFree(tmp);
	cudaThreadSynchronize();
	cudaGetLastError();
}
////////////////////////////////////////////////////////////////////////////////////////////////////
__global__ 
void __multiply(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < nRows*nCols)
	{
		c[idx].x = a[idx].x * b[idx].x;
		c[idx].y = a[idx].y * b[idx].y;
	}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
void multiply(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	float2* tmp;
	cudaMalloc((void**)&tmp, sizeof(float2)*nCols*nRows);
	
	__multiply config(nRows*nCols, 1, 256, 1, 0)	\
		(a, b, tmp, nRows, nCols);

	cudaMemcpy(c, tmp, sizeof(float2)*nCols*nRows, cudaMemcpyDeviceToDevice);
	cudaFree(tmp);
	cudaThreadSynchronize();
	cudaGetLastError();
}

//////////////////////////////////////////////////////////////////////////////////////////////////////
//__global__ 
//void __multiply(
//	float2* a,			
//	float2  b,
//	float2* c,
//	int nRows,						//Number of Rows
//	int nCols)						//Number of Cols
//{
//	int idx = threadIdx.x + blockIdx.x*blockDim.x;
//	if(idx < nRows*nCols)
//	{
//		c[idx].x = a[idx].x * b.x;
//		c[idx].y = a[idx].y * b.y;
//	}
//}
//////////////////////////////////////////////////////////////////////////////////////////////////////
//void multiply(
//	float2* a,			
//	float2  b,
//	float2* c,
//	int nRows,						//Number of Rows
//	int nCols)						//Number of Cols
//{
//	__multiply config(nRows*nCols, 1, 256, 1, 0)	\
//		(a, b, c, nRows, nCols);
//	cudaThreadSynchronize();
//	cudaGetLastError();
//}

////////////////////////////////////////////////////////////////////////////////////////////////////
__global__ 
void __divide(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < nRows*nCols)
	{
		c[idx].x = a[idx].x / b[idx].x;
		c[idx].y = a[idx].y / b[idx].y;
	}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
void divide(
	float2* a,			
	float2* b,
	float2* c,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	__divide config(nRows*nCols, 1, 256, 1, 0)	\
		(a, b, c, nRows, nCols);
	cudaThreadSynchronize();
	cudaGetLastError();
}
////////////////////////////////////////////////////////////////////////////////////////////////////
__global__ 
void __reverse(
	float2* src,			
	float2* dst,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	int idx = threadIdx.x + blockIdx.x*blockDim.x;
	if(idx < nRows*nCols)
	{
		if(src[threadIdx.x].x == 1.0f)
		{
			dst[threadIdx.x].x = 0.0f;
			dst[threadIdx.x].y = 0.0f;
		}
		else
		{
			dst[threadIdx.x].x = 1.0f;
			dst[threadIdx.x].y = 1.0f;
		}
	}
}
////////////////////////////////////////////////////////////////////////////////////////////////////
void reverse(
	float2* src,			
	float2* dst,
	int nRows,						//Number of Rows
	int nCols)						//Number of Cols
{
	__reverse config(nRows*nCols, 1, 256, 1, 0)	\
		(src, dst, nRows, nCols);
	cudaThreadSynchronize();
	cudaGetLastError();
}
}
